Skip to content

fix: Q1_0_g128 CPU dot product int truncation - #4

Closed
Marxist-Leninist wants to merge 5 commits into
PrismML-Eng:prismfrom
Marxist-Leninist:fix/q1_0_g128-cpu-kernel-int-truncation
Closed

Marxist-Leninist wants to merge 5 commits into
PrismML-Eng:prismfrom
Marxist-Leninist:fix/q1_0_g128-cpu-kernel-int-truncation

Conversation

@Marxist-Leninist

Copy link
Copy Markdown

Summary

  • The accumulator sumi in ggml_vec_dot_q1_0_g128_q8_0 was declared as int but accumulates float d1 * int sumi_block, causing the float result to be truncated to integer on each loop iteration
  • This produced garbage output for Q1_0_g128 quantized models when running on CPU (x86 and generic/portable paths)
  • Fix: change int sumi = 0 to float sumi = 0 in both kernels

Files changed

  • ggml/src/ggml-cpu/arch/x86/quants.c — x86 kernel
  • ggml/src/ggml-cpu/quants.c — generic/portable kernel

Test

Verified locally on Windows (MinGW, CPU-only) with Bonsai-8B (Q1_0_g128):

  • Before fix: garbled/random token output
  • After fix: coherent, correct responses

The accumulator `sumi` in ggml_vec_dot_q1_0_g128_q8_0 was declared
as `int` but accumulates `float d1 * int sumi_block`, causing the
float result to be truncated to integer on each iteration. This
produced garbage output for Q1_0_g128 models on CPU.

Fix: change `int sumi = 0` to `float sumi = 0` in both the x86
and generic (portable) kernels.
@github-actions github-actions Bot added the ggml label Apr 2, 2026
noreply added 4 commits April 2, 2026 16:30
The existing scalar fallback runs at ~0.2 t/s on CPUs without AVX-512
(Ryzen, Intel 12th+ gen consumer). This adds an AVX2 path that:
- Sign-extends int8->int16 in two 16-element passes per Q8_0 block
- Expands 1-bit weights to 16-bit masks via broadcast+AND+cmpeq
- Uses blendv to negate activations where weight bit=0
- Accumulates via madd_epi16 -> cvtepi32_ps -> fmadd_ps

AVX2 is supported on virtually all x86-64 CPUs from 2013+.
…sing

Replace two-pass int16 blendv approach with:
- Single-pass byte-level bit expansion (shuffle+AND+cmpeq)
- XOR+SUB negate trick (replaces slow blendv, 2-3 cyc -> 1 cyc each)
- maddubs+madd accumulation (stays in int8 longer)
- Fully unrolled k-loop (eliminates loop overhead + branch)

Benchmark on i7-10510U (AVX2+FMA, 4T):
  Scalar:  0.2 t/s prompt, 0.2 t/s gen
  AVX2 v1: 2.4 t/s prompt, 2.1 t/s gen (two-pass blendv)
  AVX2 v3: 4.7 t/s prompt, 3.1 t/s gen (this commit)

~15x faster than scalar, ~50% faster than v1.
Apply cache-blocking and prefetch optimizations from the COM6 matrix
multiplication library (github.com/Marxist-Leninist/COM6):

- Increase weight row block size from 16 to 64 for Q1_0_g128
  (1-bit rows are ~576 bytes at K=4096, 64 rows = 36KB fits in L1d)
- Add software prefetch of weight rows 4 iterations ahead,
  mirroring COM6 distributed prefetch strategy
- Enlarge tmp accumulator buffer to match larger block size

Benchmark on i7-10510U (4T, Bonsai-8B Q1_0_g128):
  Before: 3.14 t/s generation
  After:  3.43 t/s generation (+9%)
@khosravipasha

Copy link
Copy Markdown
Collaborator

This look great thanks, there was a few CPU kernel fixes and did not see them until I pushed my changes. For now removed the buggy x86, will merge one of the correct AVX ones.

Could you run the KL divergence tests described here: #8

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Fixes incorrect CPU dot-product accumulation for Q1_0_g128 quantized models by preventing float-to-int truncation during accumulation, improving correctness of inference outputs on CPU backends.

Changes:

  • Change sumi accumulator type from int to float in the generic ggml_vec_dot_q1_0_g128_q8_0 kernel.
  • Change sumi accumulator type from int to float in the x86 ggml_vec_dot_q1_0_g128_q8_0 scalar fallback.
  • Additional unrelated changes: mul_mat block-tiling/prefetch tweaks and a new .gitattributes entry for .gguf via LFS.

Reviewed changes

Copilot reviewed 4 out of 4 changed files in this pull request and generated 5 comments.

File Description
ggml/src/ggml-cpu/quants.c Fixes Q1_0_g128 x Q8_0 generic dot-product accumulation truncation by using float accumulator.
ggml/src/ggml-cpu/arch/x86/quants.c Updates Q1_0_g128 x Q8_0 accumulation type; also introduces AVX2-related changes that currently appear to break compilation.
ggml/src/ggml-cpu/ggml-cpu.c Changes tiling block size for Q1_0_g128, increases tmp buffer, and adds prefetching in matmul inner loop.
.gitattributes Adds Git LFS handling for *.gguf files.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment on lines +68 to +84
#if defined(__AVX2__)
// AVX2: single-pass byte-level processing, fully unrolled k-loop.
// Pipeline: broadcast+shuffle -> AND+cmpeq -> XOR+SUB -> maddubs+madd -> cvt+fma
const __m256i ones_8 = _mm256_set1_epi8(1);
const __m256i ones_16 = _mm256_set1_epi16(1);
const __m256i byte_shuf = _mm256_setr_epi8(
0,0,0,0,0,0,0,0, 1,1,1,1,1,1,1,1,
2,2,2,2,2,2,2,2, 3,3,3,3,3,3,3,3);
const __m256i bit_masks = _mm256_setr_epi8(
1,2,4,8,16,32,64,-128, 1,2,4,8,16,32,64,-128,
1,2,4,8,16,32,64,-128, 1,2,4,8,16,32,64,-128);
const __m256i zero = _mm256_setzero_si256();
__m256 acc = _mm256_setzero_ps();

// Unpack 32 4-bit fields into 32 bytes
// The output vector contains 32 bytes, each one in [ 0 .. 15 ] interval
static inline __m256i bytes_from_nibbles_32(const uint8_t * rsi)
{
const __m128i tmp = _mm_loadu_si128((const __m128i *)rsi);
const __m256i bytes = MM256_SET_M128I(_mm_srli_epi16(tmp, 4), tmp);
const __m256i lowMask = _mm256_set1_epi8( 0xF );
return _mm256_and_si256(lowMask, bytes);
}
for (int ib = 0; ib < nb; ++ib) {
const float d0 = GGML_CPU_FP16_TO_FP32(x[ib].d);
const uint32_t * qs32 = (const uint32_t *)x[ib].qs;

Copilot AI Apr 6, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The #if defined(__AVX2__) block here contains executable statements that reference nb, x, y_ptr, and s, but this code is located at file scope (inside the helper section) rather than inside a function. This will not compile and also appears to have overwritten the intended mul_sum_us8_pairs_float(...) implementation. Please move this logic into the appropriate vec_dot function (or restore the helper to only operate on its parameters and return a value).

Copilot uses AI. Check for mistakes.
Comment on lines +1188 to +1197
// COM6-inspired block-tiling: larger blocks for Q1_0_g128 (1-bit weights are tiny,
// so we can fit more rows in L1). Prefetch next weight block while processing current.
const int64_t blck_0 = (type == GGML_TYPE_Q1_0_g128) ? 64 : 16;
const int64_t blck_1 = 16;

const size_t src1_col_stride = src1_cont || src1->type != vec_dot_type ? row_size : nb11;

// attempt to reduce false-sharing (does not seem to make a difference)
// 16 * 2, accounting for mmla kernels
float tmp[32];
// Size: blck_0 * 2 (accounting for mmla kernels that compute 2 rows at once)
float tmp[128];

Copilot AI Apr 6, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This PR description says the fix is limited to changing int sumi -> float sumi in two dot-product kernels, but this file now includes a new block-tiling size change, a larger tmp buffer, and a new prefetch loop. Please update the PR description/title to reflect these additional behavioral/performance changes or split them into a separate PR to keep the correctness fix isolated.

Copilot uses AI. Check for mistakes.
Comment on lines +1232 to 1236
for (int64_t ir0 = iir0; ir0 < ir0_max; ir0 += num_rows_per_vec_dot) {
if (ir0 + 4 * num_rows_per_vec_dot < ir0_max) {
__builtin_prefetch(src0_row + (ir0 + 4 * num_rows_per_vec_dot) * nb01, 0, 1);
}
vec_dot(ne00, &tmp[ir0 - iir0], (num_rows_per_vec_dot > 1 ? 16 : 0), src0_row + ir0 * nb01, (num_rows_per_vec_dot > 1 ? nb01 : 0), src1_col, (num_rows_per_vec_dot > 1 ? src1_col_stride : 0), num_rows_per_vec_dot);

Copilot AI Apr 6, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

__builtin_prefetch is a GCC/Clang builtin and is not available on all supported toolchains (notably MSVC). Consider guarding this with compiler checks (e.g., __GNUC__/__clang__) and/or using an existing cross-platform prefetch abstraction (or _mm_prefetch on x86) so CPU builds remain portable.

Copilot uses AI. Check for mistakes.
Comment on lines +1196 to +1197
// Size: blck_0 * 2 (accounting for mmla kernels that compute 2 rows at once)
float tmp[128];

Copilot AI Apr 6, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

tmp is documented as blck_0 * 2, but it is hard-coded to float tmp[128] while blck_0 is now dynamic. This makes the comment inaccurate and risks future buffer overruns if blck_0 is tuned again. Prefer sizing tmp from blck_0/num_rows_per_vec_dot (e.g., VLA or max-constant with a static assert) and updating the memcpy offsets accordingly.

Suggested change
// Size: blck_0 * 2 (accounting for mmla kernels that compute 2 rows at once)
float tmp[128];
// Size: blck_0 * num_rows_per_vec_dot
const int64_t tmp_size = blck_0 * num_rows_per_vec_dot;
float tmp[tmp_size];

Copilot uses AI. Check for mistakes.
Comment on lines 668 to +672

float sumf = 0;

// Each Q1_0_g128 block has 128 elements
// Each Q8_0 block has 32 elements
// So we need 4 Q8_0 blocks per Q1_0_g128 block
#if defined(__AVX2__)
// AVX2: process 32 Q8_0 values per sub-block in two 16-element passes.

Copilot AI Apr 6, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

sumf is declared before the #if defined(__AVX2__) block, but in AVX2 builds the function assigns *s directly and never reads sumf, which will trigger an unused-variable warning. Consider moving the sumf declaration into the scalar #else path or using GGML_UNUSED(sumf) when AVX2 is enabled.

Copilot uses AI. Check for mistakes.
@khosravipasha

Copy link
Copy Markdown
Collaborator

Good new our first CPU PR just got merged int llama.cpp master branch now, if you are still working on this please rebase with PrismML's master (just pulled the main llama.cpp)

Changes: Q1_0_g128 naming is gone now, the original Q1_0 with group size 32 was deleted and Q1_0_g128 was renamed to Q1_0 now by default has group size 128.

https://github.com/PrismML-Eng/llama.cpp/tree/master

This one only has generic cpu (slow), and ARM NEON path, planning to gather the best x86 kernels from here and to send a PR there (and tag all the contributers).

@khosravipasha

Copy link
Copy Markdown
Collaborator

There is a lot of CPU PRs, planning to gether all in one and then send to the main llama.cpp
Going to close this and mention people that helped in a thread there, if you think your solution is better please comment there:
#10

bricklc referenced this pull request in bricklc/prism-ml-llama.cpp Apr 25, 2026
Complete experiment log:
  Mintplex-Labs#1  4-mag LUT:           15.1 at 8K (BEST, +38%)
  Mintplex-Labs#2  Batched extract:     13.7 (+25%)
  Mintplex-Labs#3  Inline FA block:     13.5 (I-cache pressure)
  Mintplex-Labs#4  Deferred norm:       12.9 (loses ILP)
  PrismML-Eng#5  2-pair half2:        12.0 (ternary overhead)
  PrismML-Eng#6  Select chain:        11.9 (branches kill)
  PrismML-Eng#7  Bit-arithmetic:      11.6 (ALU too heavy)
  PrismML-Eng#8  FMA branchless:      11.4 (ALU still too heavy)
  PrismML-Eng#9  Named-reg ternary:   10.3 (branches worst)
  PrismML-Eng#10 Main (8-LUT):        10.95 (baseline)
  PrismML-Eng#11 Non-vec FA:          10.2 (wrong kernel)
  Ceiling:                 24.5 (no dequant)

Apple8 hardware truth:
  1 divergent constant read < 7 ALU ops (even with fma)
  Branches cost MORE than divergent constant reads
  Array indexing ALWAYS spills on Metal
  4 constant addresses is the sweet spot

The 4-mag LUT is the dequant-level ceiling on Apple Silicon.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
bricklc referenced this pull request in bricklc/prism-ml-llama.cpp Apr 25, 2026
…boquant-kv-cache"

This reverts commit 065ef53, reversing
changes made to 7d1bd95.
retroheim referenced this pull request in retroheim/prism-ml-llama.cpp May 3, 2026
Complete experiment log:
  Mintplex-Labs#1  4-mag LUT:           15.1 at 8K (BEST, +38%)
  Mintplex-Labs#2  Batched extract:     13.7 (+25%)
  Mintplex-Labs#3  Inline FA block:     13.5 (I-cache pressure)
  Mintplex-Labs#4  Deferred norm:       12.9 (loses ILP)
  PrismML-Eng#5  2-pair half2:        12.0 (ternary overhead)
  PrismML-Eng#6  Select chain:        11.9 (branches kill)
  PrismML-Eng#7  Bit-arithmetic:      11.6 (ALU too heavy)
  PrismML-Eng#8  FMA branchless:      11.4 (ALU still too heavy)
  PrismML-Eng#9  Named-reg ternary:   10.3 (branches worst)
  PrismML-Eng#10 Main (8-LUT):        10.95 (baseline)
  PrismML-Eng#11 Non-vec FA:          10.2 (wrong kernel)
  Ceiling:                 24.5 (no dequant)

Apple8 hardware truth:
  1 divergent constant read < 7 ALU ops (even with fma)
  Branches cost MORE than divergent constant reads
  Array indexing ALWAYS spills on Metal
  4 constant addresses is the sweet spot

The 4-mag LUT is the dequant-level ceiling on Apple Silicon.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
retroheim referenced this pull request in retroheim/prism-ml-llama.cpp May 3, 2026
…kv-cache

fix: turbo4 SET_ROWS, tail-block truncation, constant coupling, stack overflow (Issue PrismML-Eng#29)
retroheim referenced this pull request in retroheim/prism-ml-llama.cpp May 3, 2026
…boquant-kv-cache"

This reverts commit 065ef53, reversing
changes made to 7d1bd95.
retroheim referenced this pull request in retroheim/prism-ml-llama.cpp May 3, 2026
Complete experiment log:
  Mintplex-Labs#1  4-mag LUT:           15.1 at 8K (BEST, +38%)
  Mintplex-Labs#2  Batched extract:     13.7 (+25%)
  Mintplex-Labs#3  Inline FA block:     13.5 (I-cache pressure)
  Mintplex-Labs#4  Deferred norm:       12.9 (loses ILP)
  PrismML-Eng#5  2-pair half2:        12.0 (ternary overhead)
  PrismML-Eng#6  Select chain:        11.9 (branches kill)
  PrismML-Eng#7  Bit-arithmetic:      11.6 (ALU too heavy)
  PrismML-Eng#8  FMA branchless:      11.4 (ALU still too heavy)
  PrismML-Eng#9  Named-reg ternary:   10.3 (branches worst)
  PrismML-Eng#10 Main (8-LUT):        10.95 (baseline)
  PrismML-Eng#11 Non-vec FA:          10.2 (wrong kernel)
  Ceiling:                 24.5 (no dequant)

Apple8 hardware truth:
  1 divergent constant read < 7 ALU ops (even with fma)
  Branches cost MORE than divergent constant reads
  Array indexing ALWAYS spills on Metal
  4 constant addresses is the sweet spot

The 4-mag LUT is the dequant-level ceiling on Apple Silicon.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
retroheim referenced this pull request in retroheim/prism-ml-llama.cpp May 3, 2026
…boquant-kv-cache"

This reverts commit 065ef53, reversing
changes made to 7d1bd95.
MasterCool4389 pushed a commit to MasterCool4389/Built_PrismML-ENG-CUDA_llama.cpp that referenced this pull request Sep 18, 2026
* Get started with Onyx

* Add architecture

* Skip keys handled in super()

* Loading tensors

* Shorten

* Graph

* Apply suggestion from @pcuenca

* Remove norm now embedding in transformers weights

* Add eot

* Explicit output_multiplier

* Handle post_norm_eps

* No super call; unhardcode eot.

The pattern `self._set_vocab_gpt2()` seems preferred throughout the
codebase, and it allows `set_vocab()` to be called from a different part
of the Python class hierarchy: the drafter model converter that we may
need eventually.

* Register for drafting

* DFlash: inherit rope type from the linked target.

Another option would be to store it in the gguf file itself.

* mmproj conversion

Note: some fields to be renamed after the implementation works. We are
keeping compatibility with the reference Meta gguf for testing purposes.

* "clip" header declarations

* Load mmproj

* Pre-processing

* Graph

* Go back to using delimiters.

Otherwise our generations are worse.

Transformers does not use them. We need to trace inputs to verify
whether they are equivalent.

* downsample_factor -> merge_size

* Add vision graph

lol, forgot from a previous commit

* Additional renames, align with llama.cpp / transformers

* Prefer _size instead of independent _h and _w

* Fix token layout

Co-authored-by: Young Han <younghan@fb.com>

* onyx: bring the chat parser onto the onyx branch

common/chat.cpp on this branch has no Onyx handling, so a converted model
serves malformed chat: the assistant preamble leaks into content
("to=self<|message|>...") and tool calls fail with

    HTTP 500 "The model produced output that does not match the expected
              peg-native format"

common_chat_params_init_onyx exists on onyx-fair-patch, added there by
8bb73dd3d. It was never on this branch, so this is not a regression --
the two lines developed independently.

The code here is taken verbatim from that commit. It is the clean side of
`git merge origin/onyx-fair-patch`: chat.cpp is one of the files that
merges without conflict. The full merge is not viable -- it produces 13
conflicts, including add/add on conversion/onyx.py and src/models/onyx.cpp
where the q_norm-folding and metadata-scale approaches contradict each
other, and PrismML-Eng#4/PrismML-Eng#7 are stacked on this branch's side of that.

Verified on this branch: builds with 0 errors, converts an Onyx checkpoint,
and serving it gives "4" for "What is 2+2?" plus a correct
get_weather {"city":"Paris"} tool call, where the unported branch gives the
two failures above.

No converter or runtime changes are included, so this should not interact
with the q_norm work.

Co-authored-by: Beto de Paola <betodepaola@meta.com>

* Less params, bilinear pos-emb interpolation as a graph op instead of CPU

* Map to symbolic V_MMPROJ instead of strings

* Make a couple params explicit

* Patchify via build_inp()

* No param for rope_theta

* Small cleanup

* Restore blank line

* Unpermute, to adapt to the latest transformers checkpoint

* Apply norm after token embeddings

This follows the latest transformers approach.

* Remove duplicated function

* build_vit

* onyx: use the model rope theta on sliding-window layers

* DFlash: conversion from transformers drafter

* Revert rope_type derivation from target

NOTE: this breaks compatibility with Meta's distributed DFlash GGUFs, as
the Q/K are stored in "NEOX" (rotated half) format, like in
transformers.

* Apply suggestion from @pcuenca

* Set model type

* Remove comment that will become obsolete

* Hardcode post_norm_rms_eps instead of new param

* Derive SWA+RoPE pattern from gguf array or scalar

* Fix model type <-> number of layers

* Reorder

* Rename

* Fix typo

* DFlash: seed the draft KV cache from multimodal embedding batches

`common_speculative_impl_draft_dflash::process()` returned early on any batch carrying embeddings, so an image prefill never had its target-layer features fused through the DFlash encoder and injected into the draft's KV cache. That left a hole spanning the image's positions, and the next injection at a post-image position failed to initialize its batch:

```
decoding image batch 1/1, n_tokens_batch = 256
decode: failed to initialize batch
llama_decode: failed to decode, ret = -1
process: llama_decode(ctx_dft) failed rc=-1 (n_tokens=17, offset=0)
srv decode: failed to process speculative batch
```

Every image request with `--spec-type draft-dflash` failed with HTTP 500. Text-only was unaffected, since those batches carry token ids and were let through.

Restore the earlier condition, which admits a batch that is either tokens or embeddings and skips only the degenerate neither/both cases. The rest of `process()` is already layout-agnostic -- it gathers features via `llama_get_embeddings_layer_inp()` and indexes `batch_in.pos[]` / `batch_in.seq_id[]`, none of which assume token ids -- so this is the whole fix.

Validated against `muse-glimmer-30B-bf16.gguf` + `mmproj-muse-glimmer-30B-bf16.gguf` + a DFlash draft head, on an image describe-the-shapes request:

- before: HTTP 500, `failed to process speculative batch`
- after: HTTP 200, draft acceptance 0.34012 (167 accepted / 491 generated), mean len 3.04

Output equivalence holds, which is the property that matters: at temperature 0 the drafted response is byte-identical to the same request served with no draft attached (1213/1213 chars), so the draft is drafting correctly through the image context rather than merely not crashing.

* Conversion: prefer rewrite to mapping

* Revert "Conversion: prefer rewrite to mapping"

This reverts commit a92d0ac.

* fix lint

* sliding_window metadata is not optional

* disable state save/load

* Apply suggestion from @pcuenca

---------

Co-authored-by: Young Han <younghan@fb.com>
Co-authored-by: Beto de Paola <betodepaola@meta.com>
Co-authored-by: Daniel Han <michaelhan2050@gmail.com>
Co-authored-by: ruanrms <ruanslv@gmail.com>
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants